#include <stdio.h>
#include <stdlib.h>
#include <string.h>
#include <time.h>

#define RED 0
#define GREEN 1
#define BLUE 2

__global__ void smooth(unsigned char *in_r, unsigned char *in_g, unsigned char *in_b, unsigned char *out_r, unsigned char *out_g, unsigned char *out_b, int width, int height){
	int k, l;
	int value = 0;
	
	int color = blockIdx.x;
	int i = threadIdx.y;
	int j = threadIdx.x;

	// Applying smooth function
	// at non-border position
	if((2 <= i) && (i < height-2) && (2 <= j) && (j < width-2)){
		if(color == RED){
			for(k=-2; k<=2; k++){
				for(l=-2; l<=2; l++){
					value += (int)in_r[(i+k)*width + (j+l)%width];
				}
			}
			out_r[i*width + j%width] = (unsigned char)value/25;
		}
		else if(color == GREEN){
			for(k=-2; k<=2; k++){
				for(l=-2; l<=2; l++){
					value += (int)in_g[(i+k)*width + (j+l)%width];
				}
			}
			out_g[i*width + j%width] = (unsigned char)value/25;
		}
		else if (color == BLUE){
			for(k=-2; k<=2; k++){
				for(l=-2; l<=2; l++){
					value += (int)in_b[(i+k)*width + (j+l)%width];
 					}
			}
			out_b[i*width + j%width] = (unsigned char)value/25;
		}
	}
}





int main(int argc, char* argv[]){
	char magic[3];
	char *input = argv[1], *output = argv[2];
	FILE *img;
	int i, j, width, height, max, value_r, value_g, value_b;
	unsigned char *mat_r, *mat_g, *mat_b, *dev_in_r, *dev_in_g, *dev_in_b, *dev_out_r, *dev_out_g, *dev_out_b;
	clock_t execution_time;



	// 1. Carregar imagem
	img = fopen(input, "rb");
	fscanf(img, "%s", magic);
	fscanf(img, "%d", &width);
	fscanf(img, "%d", &height);
	fscanf(img, "%d", &max);
	mat_r = (unsigned char*)malloc(width*height*sizeof(unsigned char));
	mat_g = (unsigned char*)malloc(width*height*sizeof(unsigned char));
	mat_b = (unsigned char*)malloc(width*height*sizeof(unsigned char));
	if(!strcmp(magic, "P6")){	// Binary code
		for(i=0; i<height; i++){
			for(j=0; j<width; j++){
				fread(&mat_r[i*width + j%width], 1, 1, img);
				fread(&mat_g[i*width + j%width], 1, 1, img);
				fread(&mat_b[i*width + j%width], 1, 1, img);
			}
		}
	}
	else if(!strcmp(magic, "P3")){	// ASCII code
		for(i=0; i<height; i++){
			for(j=0; j<width; j++){
				fscanf(img, "%d", &value_r);
				fscanf(img, "%d", &value_g);
				fscanf(img, "%d", &value_b);
				mat_r[i*width + j%width] = (unsigned char)value_r;
				mat_g[i*width + j%width] = (unsigned char)value_g;
				mat_b[i*width + j%width] = (unsigned char)value_b;
			}
		}
	}
	else{	// Erro
		printf("Erro\n");
		free(mat_r);
		free(mat_g);
		free(mat_b);
		return -1;
	}
	fclose(img);



	// 2. Copiar imagem para GPU
	cudaMalloc((void**)&dev_in_r, width*height*sizeof(unsigned char));
	cudaMalloc((void**)&dev_in_g, width*height*sizeof(unsigned char));
	cudaMalloc((void**)&dev_in_b, width*height*sizeof(unsigned char));
	cudaMalloc((void**)&dev_out_r, width*height*sizeof(unsigned char));
	cudaMalloc((void**)&dev_out_g, width*height*sizeof(unsigned char));
	cudaMalloc((void**)&dev_out_b, width*height*sizeof(unsigned char));
	cudaMemcpy(dev_in_r, mat_r, width*height*sizeof(unsigned char), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_in_g, mat_g, width*height*sizeof(unsigned char), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_in_b, mat_b, width*height*sizeof(unsigned char), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_out_r, mat_r, width*height*sizeof(unsigned char), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_out_g, mat_g, width*height*sizeof(unsigned char), cudaMemcpyHostToDevice);
	cudaMemcpy(dev_out_b, mat_b, width*height*sizeof(unsigned char), cudaMemcpyHostToDevice);



	// 3. Iniciar a contagem de tempo
	execution_time = clock();



	// 4. Processar o smooth
	dim3 threads_per_block(width, height, 1);
	smooth<<<3, threads_per_block>>>(dev_in_r, dev_in_g, dev_in_b, dev_out_r, dev_out_g, dev_out_b, width, height);



	// 5. Finalizar a contagem e exibir tempo de processamento
	execution_time = clock() - execution_time;
	printf("%.6f", ((float)execution_time/CLOCKS_PER_SEC));



	// 6. Trazer imagem processada da CPU
	cudaMemcpy(mat_r, dev_out_r, width*height*sizeof(unsigned char), cudaMemcpyDeviceToHost);
	cudaMemcpy(mat_g, dev_out_g, width*height*sizeof(unsigned char), cudaMemcpyDeviceToHost);
	cudaMemcpy(mat_b, dev_out_b, width*height*sizeof(unsigned char), cudaMemcpyDeviceToHost);



	// 7. Salvar nova imagem
	img = fopen(output, "wb");
	fprintf(img, "%s\n%d %d\n%d\n", magic, width, height, max);
	if(!strcmp(magic, "P6")){	// Binary code
		for(i=0; i<height; i++){
			for(j=0; j<width; j++){
				fwrite(&mat_r[i*width + j%width], 1, 1, img);
				fwrite(&mat_g[i*width + j%width], 1, 1, img);
				fwrite(&mat_b[i*width + j%width], 1, 1, img);
			}
		}
	}
	else if(!strcmp(magic, "P3")){	// ASCII code
		for(i=0; i<height; i++){
			for(j=0; j<width; j++){
				fprintf(img, "%d %d %d\n", mat_r[i*width + j%width], mat_g[i*width + j%width], mat_b[i*width + j%width]);
			}
		}
	}
	fclose(img);



	// 8. Liberar memória
	free(mat_r);
	free(mat_g);
	free(mat_b);
	cudaFree(dev_in_r);
	cudaFree(dev_in_g);
	cudaFree(dev_in_b);
	cudaFree(dev_out_r);
	cudaFree(dev_out_g);
	cudaFree(dev_out_b);



	// 9. FIM
	return 0;
}
